<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<title>nih: kd/cuda/knn_inline.h Source File</title>

<link href="tabs.css" rel="stylesheet" type="text/css"/>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
<link href="navtree.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="resize.js"></script>
<script type="text/javascript" src="navtree.js"></script>
<script type="text/javascript">
  $(document).ready(initResizable);
</script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
  $(document).ready(function() { searchBox.OnSelectItem(0); });
</script>

</head>
<body>
<div id="top"><!-- do not remove this div! -->


<div id="titlearea">
<table cellspacing="0" cellpadding="0">
 <tbody>
 <tr style="height: 56px;">
  
  
  <td style="padding-left: 0.5em;">
   <div id="projectname">nih
   &#160;<span id="projectnumber">0.5</span>
   </div>
   
  </td>
  
  
  
 </tr>
 </tbody>
</table>
</div>

<!-- Generated by Doxygen 1.7.5.1 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
  <div id="navrow1" class="tabs">
    <ul class="tablist">
      <li><a href="index.html"><span>Main&#160;Page</span></a></li>
      <li><a href="modules.html"><span>Modules</span></a></li>
      <li><a href="annotated.html"><span>Classes</span></a></li>
      <li class="current"><a href="files.html"><span>Files</span></a></li>
      <li>
        <div id="MSearchBox" class="MSearchBoxInactive">
        <span class="left">
          <img id="MSearchSelect" src="search/mag_sel.png"
               onmouseover="return searchBox.OnSearchSelectShow()"
               onmouseout="return searchBox.OnSearchSelectHide()"
               alt=""/>
          <input type="text" id="MSearchField" value="Search" accesskey="S"
               onfocus="searchBox.OnSearchFieldFocus(true)" 
               onblur="searchBox.OnSearchFieldFocus(false)" 
               onkeyup="searchBox.OnSearchFieldChange(event)"/>
          </span><span class="right">
            <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
          </span>
        </div>
      </li>
    </ul>
  </div>
  <div id="navrow2" class="tabs2">
    <ul class="tablist">
      <li><a href="files.html"><span>File&#160;List</span></a></li>
    </ul>
  </div>
</div>
<div id="side-nav" class="ui-resizable side-nav-resizable">
  <div id="nav-tree">
    <div id="nav-tree-contents">
    </div>
  </div>
  <div id="splitbar" style="-moz-user-select:none;" 
       class="ui-resizable-handle">
  </div>
</div>
<script type="text/javascript">
  initNavTree('knn__inline_8h.html','');
</script>
<div id="doc-content">
<div class="header">
  <div class="headertitle">
<div class="title">kd/cuda/knn_inline.h</div>  </div>
</div>
<div class="contents">
<div class="fragment"><pre class="fragment"><a name="l00001"></a>00001 <span class="comment">/*</span>
<a name="l00002"></a>00002 <span class="comment"> * Copyright (c) 2010-2011, NVIDIA Corporation</span>
<a name="l00003"></a>00003 <span class="comment"> * All rights reserved.</span>
<a name="l00004"></a>00004 <span class="comment"> *</span>
<a name="l00005"></a>00005 <span class="comment"> * Redistribution and use in source and binary forms, with or without</span>
<a name="l00006"></a>00006 <span class="comment"> * modification, are permitted provided that the following conditions are met:</span>
<a name="l00007"></a>00007 <span class="comment"> *   * Redistributions of source code must retain the above copyright</span>
<a name="l00008"></a>00008 <span class="comment"> *     notice, this list of conditions and the following disclaimer.</span>
<a name="l00009"></a>00009 <span class="comment"> *   * Redistributions in binary form must reproduce the above copyright</span>
<a name="l00010"></a>00010 <span class="comment"> *     notice, this list of conditions and the following disclaimer in the</span>
<a name="l00011"></a>00011 <span class="comment"> *     documentation and/or far materials provided with the distribution.</span>
<a name="l00012"></a>00012 <span class="comment"> *   * Neither the name of NVIDIA Corporation nor the</span>
<a name="l00013"></a>00013 <span class="comment"> *     names of its contributors may be used to endorse or promote products</span>
<a name="l00014"></a>00014 <span class="comment"> *     derived from this software without specific prior written permission.</span>
<a name="l00015"></a>00015 <span class="comment"> *</span>
<a name="l00016"></a>00016 <span class="comment"> * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS &quot;AS IS&quot; AND</span>
<a name="l00017"></a>00017 <span class="comment"> * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED</span>
<a name="l00018"></a>00018 <span class="comment"> * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE</span>
<a name="l00019"></a>00019 <span class="comment"> * DISCLAIMED. IN NO EVENT SHALL &lt;COPYRIGHT HOLDER&gt; BE LIABLE FOR ANY</span>
<a name="l00020"></a>00020 <span class="comment"> * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES</span>
<a name="l00021"></a>00021 <span class="comment"> * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;</span>
<a name="l00022"></a>00022 <span class="comment"> * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND</span>
<a name="l00023"></a>00023 <span class="comment"> * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT</span>
<a name="l00024"></a>00024 <span class="comment"> * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS</span>
<a name="l00025"></a>00025 <span class="comment"> * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.</span>
<a name="l00026"></a>00026 <span class="comment"> */</span>
<a name="l00027"></a>00027 
<a name="l00028"></a>00028 <span class="preprocessor">#include &lt;thrust/detail/backend/dereference.h&gt;</span>
<a name="l00029"></a>00029 
<a name="l00030"></a>00030 <span class="keyword">namespace </span>nih {
<a name="l00031"></a>00031 <span class="keyword">namespace </span>cuda {
<a name="l00032"></a>00032 
<a name="l00033"></a>00033 <span class="keyword">namespace </span>kd_knn {
<a name="l00034"></a>00034 
<a name="l00035"></a>00035 <span class="keyword">template</span> &lt;<span class="keyword">typename</span> VectorType, <span class="keyword">typename</span> Po<span class="keywordtype">int</span>Iterator&gt;
<a name="l00036"></a>00036 __device__ <span class="keywordtype">void</span> lookup(
<a name="l00037"></a>00037     <span class="keyword">const</span> VectorType        query,
<a name="l00038"></a>00038     <span class="keyword">const</span> Kd_node*          kd_nodes,
<a name="l00039"></a>00039     <span class="keyword">const</span> uint2*            kd_ranges,
<a name="l00040"></a>00040     <span class="keyword">const</span> uint2*            kd_leaves,
<a name="l00041"></a>00041     <span class="keyword">const</span> PointIterator     kd_points,
<a name="l00042"></a>00042     Kd_knn&lt;3&gt;::Result*      results)
<a name="l00043"></a>00043 {
<a name="l00044"></a>00044     <span class="comment">//</span>
<a name="l00045"></a>00045     <span class="comment">// 1-st pass: find the leaf containing the query point and compute an upper bound</span>
<a name="l00046"></a>00046     <span class="comment">// on the search distance</span>
<a name="l00047"></a>00047     <span class="comment">//</span>
<a name="l00048"></a>00048 
<a name="l00049"></a>00049     uint32 idx;
<a name="l00050"></a>00050     <span class="keywordtype">float</span>  dist2 = 1.0e16f;
<a name="l00051"></a>00051 
<a name="l00052"></a>00052     <span class="comment">// start from the root node</span>
<a name="l00053"></a>00053     uint32 node_index = 0;
<a name="l00054"></a>00054 
<a name="l00055"></a>00055     <span class="comment">// keep track of which leaf we visited</span>
<a name="l00056"></a>00056     uint32 first_leaf = 0;
<a name="l00057"></a>00057 
<a name="l00058"></a>00058     <span class="keywordflow">while</span> (1)
<a name="l00059"></a>00059     {
<a name="l00060"></a>00060         <span class="keyword">const</span> Kd_node node = kd_nodes[ node_index ];
<a name="l00061"></a>00061 
<a name="l00062"></a>00062         <span class="keywordflow">if</span> (node.is_leaf())
<a name="l00063"></a>00063         {
<a name="l00064"></a>00064             <span class="comment">// find the closest neighbor in this leaf</span>
<a name="l00065"></a>00065             <span class="keyword">const</span> uint2 leaf = kd_leaves[ node.get_leaf_index() ];
<a name="l00066"></a>00066             <span class="keywordflow">for</span> (uint32 i = leaf.x; i &lt; leaf.y; ++i)
<a name="l00067"></a>00067             {
<a name="l00068"></a>00068                 <span class="keyword">const</span> VectorType delta = kd_points[i] - query;
<a name="l00069"></a>00069                 <span class="keyword">const</span> <span class="keywordtype">float</span> d2 = delta[0]*delta[0] + delta[1]*delta[1] + delta[2]*delta[2];
<a name="l00070"></a>00070                 <span class="keywordflow">if</span> (dist2 &gt; d2)
<a name="l00071"></a>00071                 {
<a name="l00072"></a>00072                     dist2 = d2;
<a name="l00073"></a>00073                     idx   = i;
<a name="l00074"></a>00074                 }
<a name="l00075"></a>00075             }
<a name="l00076"></a>00076 
<a name="l00077"></a>00077             <span class="comment">// keep track of which leaf we found</span>
<a name="l00078"></a>00078             first_leaf = node_index;
<a name="l00079"></a>00079             <span class="keywordflow">break</span>;
<a name="l00080"></a>00080         }
<a name="l00081"></a>00081         <span class="keywordflow">else</span>
<a name="l00082"></a>00082         {
<a name="l00083"></a>00083             <span class="keyword">const</span> uint32 split_dim   = node.get_split_dim();
<a name="l00084"></a>00084             <span class="keyword">const</span> <span class="keywordtype">float</span>  split_plane = node.get_split_plane();
<a name="l00085"></a>00085 
<a name="l00086"></a>00086             node_index =
<a name="l00087"></a>00087                 split_dim == 0 ? node.get_child_offset() + (query[0] &lt; split_plane ? 0u : 1u) :
<a name="l00088"></a>00088                 split_dim == 1 ? node.get_child_offset() + (query[1] &lt; split_plane ? 0u : 1u) :
<a name="l00089"></a>00089                                  node.get_child_offset() + (query[2] &lt; split_plane ? 0u : 1u);
<a name="l00090"></a>00090         }
<a name="l00091"></a>00091     }
<a name="l00092"></a>00092 
<a name="l00093"></a>00093     <span class="comment">//</span>
<a name="l00094"></a>00094     <span class="comment">// 2-nd pass: visit the tree with a stack and careful pruning</span>
<a name="l00095"></a>00095     <span class="comment">//</span>
<a name="l00096"></a>00096 
<a name="l00097"></a>00097     int32  stackp = 1;
<a name="l00098"></a>00098     float4 stack[32];
<a name="l00099"></a>00099 
<a name="l00100"></a>00100     <span class="comment">// place a sentinel node in the stack</span>
<a name="l00101"></a>00101     stack[0] = make_float4( 0.0f, 0.0f, 0.0f, binary_cast&lt;float&gt;(uint32(-1)) );
<a name="l00102"></a>00102 
<a name="l00103"></a>00103     <span class="comment">// start from the root node</span>
<a name="l00104"></a>00104     node_index = 0;
<a name="l00105"></a>00105 
<a name="l00106"></a>00106     float3 cdist = make_float3( 0.0f, 0.0f, 0.0f );
<a name="l00107"></a>00107 
<a name="l00108"></a>00108     <span class="keywordflow">while</span> (node_index != uint32(-1))
<a name="l00109"></a>00109     {
<a name="l00110"></a>00110         <span class="keyword">const</span> Kd_node node = kd_nodes[ node_index ];
<a name="l00111"></a>00111 
<a name="l00112"></a>00112         <span class="keywordflow">if</span> (node.is_leaf())
<a name="l00113"></a>00113         {
<a name="l00114"></a>00114             <span class="keywordflow">if</span> (first_leaf != node_index)
<a name="l00115"></a>00115             {
<a name="l00116"></a>00116                 <span class="comment">// find the closest neighbor in this leaf</span>
<a name="l00117"></a>00117                 <span class="keyword">const</span> uint2 leaf = kd_leaves[ node.get_leaf_index() ];
<a name="l00118"></a>00118                 <span class="keywordflow">for</span> (uint32 i = leaf.x; i &lt; leaf.y; ++i)
<a name="l00119"></a>00119                 {
<a name="l00120"></a>00120                     <span class="keyword">const</span> VectorType delta = kd_points[i] - query;
<a name="l00121"></a>00121                     <span class="keyword">const</span> <span class="keywordtype">float</span> d2 = delta[0]*delta[0] + delta[1]*delta[1] + delta[2]*delta[2];
<a name="l00122"></a>00122                     <span class="keywordflow">if</span> (dist2 &gt; d2)
<a name="l00123"></a>00123                     {
<a name="l00124"></a>00124                         dist2 = d2;
<a name="l00125"></a>00125                         idx   = i;
<a name="l00126"></a>00126                     }
<a name="l00127"></a>00127                 }
<a name="l00128"></a>00128             }
<a name="l00129"></a>00129 
<a name="l00130"></a>00130             <span class="comment">// pop the next node from the stack</span>
<a name="l00131"></a>00131             <span class="keywordflow">while</span> (stackp &gt; 0)
<a name="l00132"></a>00132             {
<a name="l00133"></a>00133                 <span class="keyword">const</span> float4 stack_node = stack[ --stackp ];
<a name="l00134"></a>00134                 node_index = binary_cast&lt;uint32&gt;( stack_node.w );
<a name="l00135"></a>00135                 cdist      = make_float3( stack_node.x, stack_node.y, stack_node.z );
<a name="l00136"></a>00136 
<a name="l00137"></a>00137                 <span class="keywordflow">if</span> (cdist.x*cdist.x + cdist.y*cdist.y + cdist.z*cdist.z &gt; dist2)
<a name="l00138"></a>00138                     <span class="keywordflow">continue</span>;
<a name="l00139"></a>00139             }
<a name="l00140"></a>00140         }
<a name="l00141"></a>00141         <span class="keywordflow">else</span>
<a name="l00142"></a>00142         {
<a name="l00143"></a>00143             <span class="keyword">const</span> uint32 split_dim   = node.get_split_dim();
<a name="l00144"></a>00144             <span class="keyword">const</span> <span class="keywordtype">float</span>  split_plane = node.get_split_plane();
<a name="l00145"></a>00145 
<a name="l00146"></a>00146             <span class="keyword">const</span> <span class="keywordtype">float</span> split_dist = 
<a name="l00147"></a>00147                 split_dim == 0 ? (query[0] - split_plane) :
<a name="l00148"></a>00148                 split_dim == 1 ? (query[1] - split_plane) :
<a name="l00149"></a>00149                                  (query[2] - split_plane);
<a name="l00150"></a>00150 
<a name="l00151"></a>00151             <span class="keyword">const</span> uint32 select = split_dist &lt;= 0.0f ? 0u : 1u;
<a name="l00152"></a>00152 
<a name="l00153"></a>00153             node_index = node.get_child_offset() + select;
<a name="l00154"></a>00154 
<a name="l00155"></a>00155             <span class="comment">// compute the vector distance to the far node</span>
<a name="l00156"></a>00156             float3 cdist_far = cdist;
<a name="l00157"></a>00157             <span class="keywordflow">if</span> (split_dim == 0)      cdist_far.x = split_dist;
<a name="l00158"></a>00158             <span class="keywordflow">else</span> <span class="keywordflow">if</span> (split_dim == 1) cdist_far.y = split_dist;
<a name="l00159"></a>00159             <span class="keywordflow">else</span>                     cdist_far.z = split_dist;
<a name="l00160"></a>00160 
<a name="l00161"></a>00161             <span class="comment">// check whether we should push the far node on the stack</span>
<a name="l00162"></a>00162             <span class="keyword">const</span> <span class="keywordtype">float</span> dist_far2 = cdist_far.x*cdist_far.x +
<a name="l00163"></a>00163                                     cdist_far.y*cdist_far.y +
<a name="l00164"></a>00164                                     cdist_far.z*cdist_far.z;
<a name="l00165"></a>00165 
<a name="l00166"></a>00166             <span class="keywordflow">if</span> (dist_far2 &lt;= dist2)
<a name="l00167"></a>00167             {
<a name="l00168"></a>00168                 stack[ stackp++ ] = make_float4(
<a name="l00169"></a>00169                     cdist_far.x,
<a name="l00170"></a>00170                     cdist_far.y,
<a name="l00171"></a>00171                     cdist_far.z,
<a name="l00172"></a>00172                     binary_cast&lt;float&gt;( node.get_child_offset() + 1u - select ) );
<a name="l00173"></a>00173             }
<a name="l00174"></a>00174         }
<a name="l00175"></a>00175     }
<a name="l00176"></a>00176 
<a name="l00177"></a>00177     <span class="comment">// write the result</span>
<a name="l00178"></a>00178     results-&gt;index = idx;
<a name="l00179"></a>00179     results-&gt;dist2 = dist2;
<a name="l00180"></a>00180 }
<a name="l00181"></a>00181 
<a name="l00182"></a>00182 <span class="keyword">template</span> &lt;<span class="keyword">typename</span> QueryIterator, <span class="keyword">typename</span> Po<span class="keywordtype">int</span>Iterator&gt;
<a name="l00183"></a>00183 __global__ <span class="keywordtype">void</span> lookup_kernel(
<a name="l00184"></a>00184     <span class="keyword">const</span> uint32                    n_points,
<a name="l00185"></a>00185     <span class="keyword">const</span> QueryIterator             points_begin,
<a name="l00186"></a>00186     <span class="keyword">const</span> Kd_node*                  kd_nodes,
<a name="l00187"></a>00187     <span class="keyword">const</span> uint2*                    kd_ranges,
<a name="l00188"></a>00188     <span class="keyword">const</span> uint2*                    kd_leaves,
<a name="l00189"></a>00189     <span class="keyword">const</span> PointIterator             kd_points,
<a name="l00190"></a>00190     Kd_knn&lt;3&gt;::Result*              results)
<a name="l00191"></a>00191 {
<a name="l00192"></a>00192     <span class="keyword">const</span> uint32 grid_size = gridDim.x * blockDim.x;
<a name="l00193"></a>00193 
<a name="l00194"></a>00194     <span class="comment">// loop through all logical blocks associated to this physical one</span>
<a name="l00195"></a>00195     <span class="keywordflow">for</span> (uint32 base_idx = blockIdx.x * blockDim.x;
<a name="l00196"></a>00196                 base_idx &lt; n_points;
<a name="l00197"></a>00197                 base_idx += grid_size)
<a name="l00198"></a>00198     {
<a name="l00199"></a>00199         <span class="keyword">const</span> uint32 index = threadIdx.x + base_idx;
<a name="l00200"></a>00200         <span class="keywordflow">if</span> (index &gt;= n_points)
<a name="l00201"></a>00201             <span class="keywordflow">return</span>;
<a name="l00202"></a>00202 
<a name="l00203"></a>00203         lookup(
<a name="l00204"></a>00204             thrust::detail::backend::dereference( points_begin + index ),
<a name="l00205"></a>00205             kd_nodes,
<a name="l00206"></a>00206             kd_ranges,
<a name="l00207"></a>00207             kd_leaves,
<a name="l00208"></a>00208             kd_points,
<a name="l00209"></a>00209             results + index );
<a name="l00210"></a>00210     }
<a name="l00211"></a>00211 }
<a name="l00212"></a>00212 
<a name="l00213"></a>00213 } <span class="comment">// namespace knn</span>
<a name="l00214"></a>00214 
<a name="l00215"></a>00215 <span class="comment">// perform a k-nn lookup for a set of query points</span>
<a name="l00216"></a>00216 <span class="comment">//</span>
<a name="l00217"></a>00217 <span class="comment">// \param points_begin     beginning of the query point sequence</span>
<a name="l00218"></a>00218 <span class="comment">// \param points_end       end of the query point sequence</span>
<a name="l00219"></a>00219 <span class="comment">// \param kd_nodes         k-d tree nodes</span>
<a name="l00220"></a>00220 <span class="comment">// \param kd_ranges        k-d tree node ranges</span>
<a name="l00221"></a>00221 <span class="comment">// \param kd_leaves        k-d tree leaves</span>
<a name="l00222"></a>00222 <span class="comment">// \param kd_points        k-d tree points</span>
<a name="l00223"></a>00223 <span class="keyword">template</span> &lt;<span class="keyword">typename</span> QueryIterator, <span class="keyword">typename</span> Po<span class="keywordtype">int</span>Iterator&gt;
<a name="l00224"></a><a class="code" href="structnih_1_1cuda_1_1_kd__knn_3_013_01_4.html#a70c1237dd1598e801322c991705f8e35">00224</a> <span class="keywordtype">void</span> <a class="code" href="structnih_1_1cuda_1_1_kd__knn.html">Kd_knn&lt;3&gt;::run</a>(
<a name="l00225"></a>00225     <span class="keyword">const</span> QueryIterator             points_begin,
<a name="l00226"></a>00226     <span class="keyword">const</span> QueryIterator             points_end,
<a name="l00227"></a>00227     <span class="keyword">const</span> <a class="code" href="structnih_1_1_kd__node.html">Kd_node</a>*                  kd_nodes,
<a name="l00228"></a>00228     <span class="keyword">const</span> uint2*                    kd_ranges,
<a name="l00229"></a>00229     <span class="keyword">const</span> uint2*                    kd_leaves,
<a name="l00230"></a>00230     <span class="keyword">const</span> PointIterator             kd_points,
<a name="l00231"></a>00231     Result*                         results)
<a name="l00232"></a>00232 {
<a name="l00233"></a>00233     <span class="keyword">const</span> uint32 n_points = uint32( points_end - points_begin );
<a name="l00234"></a>00234 
<a name="l00235"></a>00235     <span class="keyword">const</span> uint32 BLOCK_SIZE = 128;
<a name="l00236"></a>00236     <span class="keyword">const</span> <span class="keywordtype">size_t</span> max_blocks = thrust::detail::backend::cuda::arch::max_active_blocks(
<a name="l00237"></a>00237         kd_knn::lookup_kernel&lt;QueryIterator,PointIterator&gt;, BLOCK_SIZE, 0);
<a name="l00238"></a>00238     <span class="keyword">const</span> <span class="keywordtype">size_t</span> n_blocks   = <a class="code" href="group__vectors.html#ga5809911d15c32de18955c59865379695">nih::min</a>( max_blocks, (n_points + BLOCK_SIZE-1) / BLOCK_SIZE );
<a name="l00239"></a>00239 
<a name="l00240"></a>00240     kd_knn::lookup_kernel&lt;&lt;&lt;n_blocks,BLOCK_SIZE&gt;&gt;&gt; (
<a name="l00241"></a>00241         n_points,
<a name="l00242"></a>00242         points_begin,
<a name="l00243"></a>00243         kd_nodes,
<a name="l00244"></a>00244         kd_ranges,
<a name="l00245"></a>00245         kd_leaves,
<a name="l00246"></a>00246         kd_points,
<a name="l00247"></a>00247         results );
<a name="l00248"></a>00248 
<a name="l00249"></a>00249     cudaThreadSynchronize();
<a name="l00250"></a>00250 }
<a name="l00251"></a>00251 
<a name="l00252"></a>00252 } <span class="comment">// namespace cuda</span>
<a name="l00253"></a>00253 } <span class="comment">// namespace nih</span>
<a name="l00254"></a>00254 
</pre></div></div>
</div>
  <div id="nav-path" class="navpath">
    <ul>
      <li class="navelem"><b>knn_inline.h</b>      </li>
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
     onmouseover="return searchBox.OnSearchSelectShow()"
     onmouseout="return searchBox.OnSearchSelectHide()"
     onkeydown="return searchBox.OnSearchSelectKey(event)">
<a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(0)"><span class="SelectionMark">&#160;</span>All</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(1)"><span class="SelectionMark">&#160;</span>Classes</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(2)"><span class="SelectionMark">&#160;</span>Files</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(3)"><span class="SelectionMark">&#160;</span>Functions</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(4)"><span class="SelectionMark">&#160;</span>Variables</a></div>

<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0" 
        name="MSearchResults" id="MSearchResults">
</iframe>
</div>


    <li class="footer">Generated on Thu Dec 15 2011 11:56:19 for nih by
    <a href="http://www.doxygen.org/index.html">
    <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.7.5.1 </li>
   </ul>
 </div>


</body>
</html>
